#include "sat_cl.h"
#include "scan.h"
#include "transpose.h"
#include "hr_time.h"

//#define DEBUG 

#ifdef DEBUG
void print_arr(SAT_test_t * t, cl_float4 *arr);
#endif

// got a clever implementation of this courtesy of: http://www.gamedev.net/topic/229831-nearest-power-of-2/
unsigned NextPow2( unsigned n )
{
	const int MantissaMask = (1<<23) - 1;

	(*(float*)&n) = (float) n;
	n = n + MantissaMask & ~MantissaMask;
	n = (unsigned) (*(float*)&n);

	return n;
} 

// SAT scan helper
cl_mem sat_scan_helper(SAT_test_t *t, const unsigned int width, const unsigned int height, 
					 unsigned int stride, size_t numBlocks, size_t blockSize)
{
	size_t szLocalWorkSize[2];
	size_t szGlobalWorkSize[2];
	cl_int errcode_ret;
	
	szLocalWorkSize[0] = blockSize;
	szLocalWorkSize[1] = 1;
    szGlobalWorkSize[0] = width;//numBlocks*blockSize;
    szGlobalWorkSize[1] = height;

	// setup parameter values 
	clSetKernelArg(t->hSATScanKernel, 0, sizeof(cl_mem), (void*)&(t->dev_src)); 
	clSetKernelArg(t->hSATScanKernel, 1, sizeof(cl_mem), (void*)&(t->dev_dst)); 
	clSetKernelArg(t->hSATScanKernel, 2, sizeof(cl_uint), (void*)&width);
	clSetKernelArg(t->hSATScanKernel, 3, sizeof(cl_uint), (void*)&stride);
	errcode_ret = clEnqueueNDRangeKernel(t->hCmdQueue, t->hSATScanKernel, 2, NULL, szGlobalWorkSize, szLocalWorkSize,
		0, NULL, NULL);
	ERR_CHECK(errcode_ret, "clEnqueueNDRangeKernel");

	// need an additional iteration of not a power of 2
	if (stride == NextPow2((unsigned int)ceil((float)width / 2))) {
		return t->dev_dst;
	}
	else {
		// ping pong buffers
		cl_mem temp = t->dev_src;
		t->dev_src = t->dev_dst;
		t->dev_dst = temp;
		return sat_scan_helper(t, width, height, stride << 1, numBlocks, blockSize);
	}
}

cl_mem sat_scan(SAT_test_t *t, const unsigned int width, const unsigned int height) {
	size_t blockSize = t->blockSize;
	size_t numBlocks = (unsigned int) ceil( (float) width / blockSize );

	return sat_scan_helper(t, width, height, 1, numBlocks, blockSize);
}


void launch_naive_transpose(SAT_test_t *t, cl_mem dev_i_data, 
        cl_mem dev_o_data,
        int rows, int cols)
{
   //NAIVE MATRIX TRANPOSE 
    //rows of dev_i_data into columns of dev_o_data
	size_t szLocalWorkSize[2];
	size_t szGlobalWorkSize[2];
	
	szLocalWorkSize[0] = 8;
	szLocalWorkSize[1] = 8;
    szGlobalWorkSize[0] = cols;
    szGlobalWorkSize[1] = rows;

	// setup parameter values 
	clSetKernelArg(t->hTransposeKernel, 0, sizeof(cl_mem), (void*)&dev_i_data); 
	clSetKernelArg(t->hTransposeKernel, 1, sizeof(cl_mem), (void*)&dev_o_data); 
	clSetKernelArg(t->hTransposeKernel, 2, sizeof(cl_int), (void*)&rows);
	clSetKernelArg(t->hTransposeKernel, 3, sizeof(cl_int), (void*)&cols);
	clEnqueueNDRangeKernel(t->hCmdQueue, t->hTransposeKernel, 2, NULL, szGlobalWorkSize, szLocalWorkSize,
		0, NULL, NULL);
}

enum Run {Scan, Transpose};

void split_time(SAT_test_t *t, Run r, CStopWatch * s) {
		clFlush(t->hCmdQueue);
		clFinish(t->hCmdQueue);
		s->stopTimer();
		if (r == Scan)
			t->scan_time += s->getElapsedTime();
		else if (r == Transpose)
			t->transpose_time += s->getElapsedTime();
}

double launch_SAT(SAT_test_t *t)
{
		float time = 0.0f;
		cl_int errcode_ret;

		printf("Copying input matrix to GPU...\n");
		errcode_ret = clEnqueueWriteBuffer(t->hCmdQueue, t->dev_src, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
		t->hst_src, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueWriteBuffer dev_src");

		CStopWatch tim;

#ifdef DEBUG
		cl_float4 * tempa = (cl_float4 *) malloc (t->num * t->num * sizeof(cl_float4));
		errcode_ret = clEnqueueReadBuffer(t->hCmdQueue, t->dev_src, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
			tempa, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueReadBuffer");
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s1;
		s1.startTimer();
		cl_mem temp;
		tim.startTimer();
		temp = sat_scan(t, t->num, t->num);
		if (temp == t->dev_src) {
			t->dev_src = t->dev_dst;
			t->dev_dst = temp;
		}
		if (t->split) {
			split_time(t, Scan, &s1);
		}


#ifdef DEBUG
		errcode_ret = clEnqueueReadBuffer(t->hCmdQueue, t->dev_dst, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
		tempa, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueReadBuffer");
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s2;
		s2.startTimer();
		launch_naive_transpose(t, t->dev_dst, t->dev_src, t->num, t->num);
		if (t->split) {
			split_time(t, Transpose, &s2);
		}

#ifdef DEBUG
		errcode_ret = clEnqueueReadBuffer(t->hCmdQueue, t->dev_src, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
		tempa, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueReadBuffer");
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s3;
		s3.startTimer();
		temp = sat_scan(t, t->num, t->num);
		if (temp == t->dev_src) {
			t->dev_src = t->dev_dst;
			t->dev_dst = temp;
		}
		if (t->split) {
			split_time(t, Scan, &s3);
		}

#ifdef DEBUG
		errcode_ret = clEnqueueReadBuffer(t->hCmdQueue, t->dev_dst, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
		tempa, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueReadBuffer");
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s4;
		s4.startTimer();
		launch_naive_transpose(t, t->dev_dst, t->dev_src, t->num, t->num);
		temp = t->dev_dst;
		t->dev_dst = t->dev_src;
		t->dev_src = temp;
		if (t->split) {
			split_time(t, Transpose, &s4);
		}
		
		clFlush(t->hCmdQueue);
		clFinish(t->hCmdQueue);
		tim.stopTimer();

#ifdef DEBUG
		errcode_ret = clEnqueueReadBuffer(t->hCmdQueue, t->dev_dst, CL_TRUE, 0, t->num * t->num * sizeof(cl_float4),
		tempa, 0, NULL, NULL);
		ERR_CHECK(errcode_ret, "clEnqueueReadBuffer");
		print_arr(t, tempa);
		free(tempa);
		printf("\n");
#endif

		return tim.getElapsedTime();
}
